#ifndef __OPERATION_H
#define __OPERATION_H

#include "tensor.h"
#include "utils.h"

namespace mslite {

template <typename Dtype>
class Operation {
public:
    Operation() = default;
    Operation(std::string name);
    virtual ~Operation();

    // Prepare
    inline void setLocal(std::vector<size_t> local) {local_ = local;}
    inline void setGlobal(std::vector<size_t> global) {global_ = global;}
    virtual int bind(void) = 0;

    // Run
    int run(cl_event &profEvent);
    virtual int setInputs(const std::vector<std::vector<Dtype>>& inputs) = 0;
    virtual std::vector<std::vector<Dtype>> getOutput(void) = 0;

protected:
    cl_program program_;
    cl_kernel kernel_;
    std::string name_;
    std::string code_;
    std::vector<size_t> global_;
    std::vector<size_t> local_;
    std::vector<Tensor<Dtype>> inputs_;
    std::vector<Tensor<Dtype>> outputs_;
private:
    int compile(void);
};

template<typename Dtype>
Operation<Dtype>::Operation(std::string name) {
    name_ = name;
    readSrcCode(name_ + ".cl", code_);
    compile();
}

template<typename Dtype>
Operation<Dtype>::~Operation() {
    cl_int status;
    for (auto tensor: inputs_)
        if (tensor.eleNums() != 0)
            status = clReleaseMemObject(tensor.memObj_);
    for (auto tensor: outputs_)
        if (tensor.eleNums() != 0)
            status = clReleaseMemObject(tensor.memObj_);
    status = clReleaseKernel(kernel_);
    status = clReleaseProgram(program_);
}

template<typename Dtype>
int Operation<Dtype>::compile(void) {
    cl_int status;
    const char *srcCode_[] = {code_.c_str()};
    size_t srcSize = code_.size();
    const size_t *srcSize_ = {&srcSize};
    program_ = clCreateProgramWithSource(Enviroment::context_, 1, srcCode_, srcSize_, &status); 
    CHECK(status);
    status = clBuildProgram(program_, 1, &Enviroment::device_, NULL, NULL, NULL);
    CHECK(status);
    kernel_ = clCreateKernel(program_, name_.c_str(), &status);  // 得到指定名字的内核实例的句柄
    CHECK(status);
    return status;
}

template<typename Dtype>
int Operation<Dtype>::run(cl_event &profEvent) {
    cl_int status;
    status = clEnqueueNDRangeKernel(Enviroment::queue_, kernel_, global_.size(), NULL, global_.data(), local_.data(), 0, NULL, &profEvent);
    CHECK(status);
    clWaitForEvents(1, &profEvent);
    status = clFinish(Enviroment::queue_);
    return status;
}

}

#endif